#ifndef _CUDA_MAXPOOLING_CU_
#define _CUDA_MAXPOOLING_CU_
namespace cuda_matrix{
    namespace __maxpooling{
       
       
    };

    /* 
       pool src to dst,optimized for pool_size = 2^pool_bits
       with block shape < 16 , 16 >
    */
    temlate<int map_method, int st_method, int pool_bits>
    __device__ void __pool_procedure_1616( int block_x,
                                           int block_y,    
                                           float s_mm[16][17],
                                           Mat dst,
                                           const Mat src  ){
        // load from src 
        for( int yy = 0 ; yy < (16<<pool_bits)  ; yy += 16 ){
            for( int xx = 0 ; xx < (16<<pool_bits) ; xx += 16 ){                

                int x_idx = block_x * (16 << pool_bits) + xx + threadIdx.x;
                int y_idx = block_y * (16 << pool_bits) + yy + threadIdx.y; 
                // load data into memory 
                if( x_idx < src.x_max && y_idx < src.y_max ) {
                    s_mm[ threadIdx.y ][ threadIdx.x ] = src[ y_idx ][ x_idx ];
                }
                __syncthreads();
                
                //cuda_reduce::reduce_block_1616<>()
                

            }
        }
    }
};

#endif
